

void f_ConvolutionNet_2(__global float* nodeVar, __global float* newNodeVar, __global float* WeightsNode, uint2* weightsOffset, ushort layer);
void f_ConvolutionNet_3(__global float* nodeVar, __global float* newNodeVar, __global float* WeightsNode, uint2* weightsOffset, ushort layer);
void f_ConvolutionNet_4(__global float* nodeVar, __global float* newNodeVar, __global float* WeightsNode, uint2* weightsOffset, ushort layer);
void f_ConvolutionNet_8(__global float* nodeVar, __global float* newNodeVar, __global float* WeightsNode, uint2* weightsOffset, ushort layer);
void f_ConvolutionNet_16(__global float* nodeVar, __global float* newNodeVar, __global float* WeightsNode, uint2* weightsOffset, ushort layer);


//typedef void(^func_SamplingBlock)(void);

__kernel void f_aiInputLayerFeedforward(__global float* inData,
                                        __global float* nodeVar,
                                        __global float* newNodeVar,
                                        __global float* WeightsNode, uint* weightsOffset, uchar BitType) {
    //size_t id = get_group_id(1) * get_num_group(1) + get_group_id(0);
    size_t id = get_global_linear_id();
    if(id >= weightsOffset[0].s0) return;
    
    
    nodeVar[id] = (1.0 / exp(inData[id] * WeightsNode[id]));
    printf("weights %f = ", WeightsNode[id]);
    return;
    if(id == 0){
        size_t work_size[2];
        
        work_size[0] = sqrt((float)weightsOffset[1].s1)+1;
        work_size[1] = work_size[0];
        ndrange_t ndrange1 = ndrange_2D(work_size);
        
        int err_ret = 0;
        
        if(BitType == 2){
            void(^func)(void) = ^{ f_ConvolutionNet_2(nodeVar, newNodeVar, WeightsNode, weightsOffset, 1); };
            err_ret = enqueue_kernel(get_default_queue(), CLK_ENQUEUE_FLAGS_WAIT_KERNEL, ndrange1, func);
        } else if(BitType == 3){
            void(^func)(void) = ^{ f_ConvolutionNet_3(nodeVar, newNodeVar, WeightsNode, weightsOffset, 1); };
            err_ret = enqueue_kernel(get_default_queue(), CLK_ENQUEUE_FLAGS_WAIT_KERNEL, ndrange1, func);
        } else if(BitType == 4){
            void(^func)(void) = ^{ f_ConvolutionNet_4(nodeVar, newNodeVar, WeightsNode, weightsOffset, 1); };
            err_ret = enqueue_kernel(get_default_queue(), CLK_ENQUEUE_FLAGS_WAIT_KERNEL, ndrange1, func);
        } else if(BitType == 8){
            void(^func)(void) = ^{ f_ConvolutionNet_8(nodeVar, newNodeVar, WeightsNode, weightsOffset, 1); };
            err_ret = enqueue_kernel(get_default_queue(), CLK_ENQUEUE_FLAGS_WAIT_KERNEL, ndrange1, func);
        } else if(BitType == 16){
            void(^func)(void) = ^{ f_ConvolutionNet_16(nodeVar, newNodeVar, WeightsNode, weightsOffset, 1); };
            err_ret = enqueue_kernel(get_default_queue(), CLK_ENQUEUE_FLAGS_WAIT_KERNEL, ndrange1, func);
        }

    }
    
    
}


#define DEF_FUNC(bit_) \
    if(id == 0 && weightsOffset[layer+1]){ \
        size_t work_size[2]; \
        \
        work_size[0] = sqrt((float)(weightsOffset[layer+1]))+1; \
        work_size[1] = work_size[0]; \
        ndrange_t ndrange1 = ndrange_2D(work_size); \
        \
        if(layer&1){ \
            void(^func_SamplingBlock)(void) = ^{ f_ConvolutionNet_#bit_(newNodeVar, nodeVar, WeightsNode, weightsOffset, layer+1); }; \
            int err_ret = enqueue_kernel(get_default_queue(), CLK_ENQUEUE_FLAGS_WAIT_KERNEL, ndrange1, func_SamplingBlock); \
        } else { \
            void(^func_SamplingBlock)(void) = ^{ f_ConvolutionNet_#bit_(nodeVar, WeightsNode, newNodeVar, weightsOffset, layer+1); }; \
            int err_ret = enqueue_kernel(get_default_queue(), CLK_ENQUEUE_FLAGS_WAIT_KERNEL, ndrange1, func_SamplingBlock); \
        } \
    }\



//调用卷积函数


//卷积
void f_ConvolutionNet_2(__global float* nodeVar, __global float* newNodeVar, __global float* WeightsNode, uint* weightsOffset, ushort layer) {
    size_t id = get_global_linear_id();
    size_t offset = weightsOffset[layer].s0*2;
    if(id >= weightsOffset[layer].s1) return;
    //float gnewNodeVar = 0.0;
    
    size_t offset_id = id*2 * 2;
    float2 v1 = vload2(offset_id, nodeVar) * WeightsNode[offset_id + offset];
    v1 += vload2(offset_id+2, nodeVar) * WeightsNode[(offset_id+2) + offset];
    v1 += vload2(offset_id+4, nodeVar) * WeightsNode[(offset_id+2) + offset];
    
    vstore2(v1, id*2, newNodeVar);
    
    if(id == 0 && weightsOffset[layer+1]){ 
        size_t work_size[2]; 
        
        work_size[0] = sqrt((float)(weightsOffset[layer+1]))+1; 
        work_size[1] = work_size[0]; 
        ndrange_t ndrange1 = ndrange_2D(work_size); 
        
        if(layer&1){ 
            void(^func_SamplingBlock)(void) = ^{ f_ConvolutionNet_2(newNodeVar, nodeVar, WeightsNode, weightsOffset, layer+1); } ;
            int err_ret = enqueue_kernel(get_default_queue(), CLK_ENQUEUE_FLAGS_WAIT_KERNEL, ndrange1, func_SamplingBlock); 
        } else { 
            void(^func_SamplingBlock)(void) = ^{ f_ConvolutionNet_2(nodeVar, WeightsNode, newNodeVar, weightsOffset, layer+1); } ;
            int err_ret = enqueue_kernel(get_default_queue(), CLK_ENQUEUE_FLAGS_WAIT_KERNEL, ndrange1, func_SamplingBlock); 
        } 
    }
}

void f_ConvolutionNet_3(__global float* nodeVar, __global float* newNodeVar, __global float* WeightsNode, uint* weightsOffset, ushort layer) {
    size_t id = get_global_linear_id();
    size_t offset = weightsOffset[layer];
    if(id >= weightsOffset[layer]) return;
    //float gnewNodeVar = 0.0;
    
    size_t offset_id = id*2 * 3;
    float3 v1 = vload3(offset_id, nodeVar) * WeightsNode[offset_id + offset];
    v1 += vload3(offset_id+3, nodeVar) * WeightsNode[(offset_id+3) + offset];
    v1 += vload3(offset_id+6, nodeVar) * WeightsNode[(offset_id+6) + offset];
    
    vstore3(v1, id*3, newNodeVar);
    
    if(id == 0 && weightsOffset[layer+1]){ 
        size_t work_size[2]; 
        
        work_size[0] = sqrt((float)(weightsOffset[layer+1]))+1; 
        work_size[1] = work_size[0]; 
        ndrange_t ndrange1 = ndrange_2D(work_size); 
        
        if(layer&1){ 
            void(^func_SamplingBlock)(void) = ^{ f_ConvolutionNet_3(newNodeVar, nodeVar, WeightsNode, weightsOffset, layer+1); } ;
            int err_ret = enqueue_kernel(get_default_queue(), CLK_ENQUEUE_FLAGS_WAIT_KERNEL, ndrange1, func_SamplingBlock); 
        } else { 
            void(^func_SamplingBlock)(void) = ^{ f_ConvolutionNet_3(nodeVar, WeightsNode, newNodeVar, weightsOffset, layer+1); } ;
            int err_ret = enqueue_kernel(get_default_queue(), CLK_ENQUEUE_FLAGS_WAIT_KERNEL, ndrange1, func_SamplingBlock); 
        } 
    }
}

void f_ConvolutionNet_4(__global float* nodeVar, __global float* newNodeVar, __global float* WeightsNode, uint* weightsOffset, ushort layer) {
    size_t id = get_global_linear_id();
    size_t offset = weightsOffset[layer];
    if(id >= weightsOffset[layer]) return;
    
    size_t offset_id = id*2 * 4;
    float4 v1 = vload4(offset_id, nodeVar) * WeightsNode[offset_id + offset];
    v1 += vload4(offset_id+4, nodeVar) * WeightsNode[(offset_id+4) + offset];
    v1 += vload4(offset_id+8, nodeVar) * WeightsNode[(offset_id+8) + offset];
    
    vstore4(v1, id*4, newNodeVar);
    
    if(id == 0 && weightsOffset[layer+1]){ 
        size_t work_size[2]; 
        
        work_size[0] = sqrt((float)(weightsOffset[layer+1]))+1; 
        work_size[1] = work_size[0]; 
        ndrange_t ndrange1 = ndrange_2D(work_size); 
        
        if(layer&1){ 
            void(^func_SamplingBlock)(void) = ^{ f_ConvolutionNet_4(newNodeVar, nodeVar, WeightsNode, weightsOffset, layer+1); } ;
            int err_ret = enqueue_kernel(get_default_queue(), CLK_ENQUEUE_FLAGS_WAIT_KERNEL, ndrange1, func_SamplingBlock); 
        } else { 
            void(^func_SamplingBlock)(void) = ^{ f_ConvolutionNet_4(nodeVar, WeightsNode, newNodeVar, weightsOffset, layer+1); } ;
            int err_ret = enqueue_kernel(get_default_queue(), CLK_ENQUEUE_FLAGS_WAIT_KERNEL, ndrange1, func_SamplingBlock); 
        } 
    }
}

void f_ConvolutionNet_8(__global float* nodeVar, __global float* newNodeVar, __global float* WeightsNode, uint* weightsOffset, ushort layer) {
    size_t id = get_global_linear_id();
    size_t offset = weightsOffset[layer];
    if(id >= offset) return;
    
    size_t offset_id = id*2 * 8;
    float8 v1 = vload8(offset_id, nodeVar) * WeightsNode[offset_id  + offset];
    v1 += vload8(offset_id+8, nodeVar) * WeightsNode[(offset_id+8) + offset];
    v1 += vload8(offset_id+16, nodeVar) * WeightsNode[(offset_id+16) + offset];
    
    vstore8(v1, id*8, newNodeVar);
    
    if(id == 0 && weightsOffset[layer+1]){ 
        size_t work_size[2]; 
        
        work_size[0] = sqrt((float)(weightsOffset[layer+1]))+1; 
        work_size[1] = work_size[0]; 
        ndrange_t ndrange1 = ndrange_2D(work_size); 
        
        if(layer&1){ 
            void(^func_SamplingBlock)(void) = ^{ f_ConvolutionNet_8(newNodeVar, nodeVar, WeightsNode, weightsOffset, layer+1); } ;
            int err_ret = enqueue_kernel(get_default_queue(), CLK_ENQUEUE_FLAGS_WAIT_KERNEL, ndrange1, func_SamplingBlock); 
        } else { 
            void(^func_SamplingBlock)(void) = ^{ f_ConvolutionNet_8(nodeVar, WeightsNode, newNodeVar, weightsOffset, layer+1); } ;
            int err_ret = enqueue_kernel(get_default_queue(), CLK_ENQUEUE_FLAGS_WAIT_KERNEL, ndrange1, func_SamplingBlock); 
        } 
    }
}

void f_ConvolutionNet_16(__global float* nodeVar, __global float* WeightsNode, __global float* newNodeVar, uint* weightsOffset, ushort layer) {
    size_t id = get_global_linear_id();
    size_t offset = weightsOffset[layer];
    
    //float gnewNodeVar = 0.0;
    
    size_t offset_id = id*2 * 16;
    float16 v1 = vload16(offset_id, nodeVar) * WeightsNode[offset_id  + offset];
    v1 += vload16(offset_id+16, nodeVar) * WeightsNode[(offset_id+16) + offset];
    v1 += vload16(offset_id+32, nodeVar) * WeightsNode[(offset_id+32) + offset];
    
    //gnewNodeVar = v1.s0 + v1.s1 + v1.s2 + v1.s3 + v1.s4 + v1.s5 + v1.s6 + v1.s7 + v1.s8 + v1.s9 + v1.sa + v1.sb + v1.sc + v1.sd + v1.se + v1.sf;
    //newNodeVar[id] = gnewNodeVar;
    vstore16(v1, id*16, newNodeVar);
    
    if(id == 0 && weightsOffset[layer+1]){ 
        size_t work_size[2]; 
        
        work_size[0] = sqrt((float)(weightsOffset[layer+1]))+1; 
        work_size[1] = work_size[0]; 
        ndrange_t ndrange1 = ndrange_2D(work_size); 
        
        if(layer&1){ 
            void(^func_SamplingBlock)(void) = ^{ f_ConvolutionNet_16(newNodeVar, nodeVar, WeightsNode, weightsOffset, layer+1); } ;
            int err_ret = enqueue_kernel(get_default_queue(), CLK_ENQUEUE_FLAGS_WAIT_KERNEL, ndrange1, func_SamplingBlock); 
        } else { 
            void(^func_SamplingBlock)(void) = ^{ f_ConvolutionNet_16(nodeVar, WeightsNode, newNodeVar, weightsOffset, layer+1); } ;
            int err_ret = enqueue_kernel(get_default_queue(), CLK_ENQUEUE_FLAGS_WAIT_KERNEL, ndrange1, func_SamplingBlock); 
        } 
    }
}





